#ifndef _CG_INFO_H_
#define _CG_INFO_H_
/*
** Define: _CG_VERSION
*/
#if defined(__MACACC__)
#include "__clang_maca_device_functions.h"
#endif

#define _CG_VERSION 1000

/*
** Define: _CG_ABI_VERSION
*/
#ifndef _CG_ABI_VERSION
#define _CG_ABI_VERSION 1
#endif

/*
** Define: _CG_ABI_EXPERIMENTAL
** Desc: If enabled, sets all features enabled (ABI-breaking or experimental)
*/
#if defined(_CG_ABI_EXPERIMENTAL)
#endif

#define _CG_CONCAT_INNER(x, y) x##y
#define _CG_CONCAT_OUTER(x, y) _CG_CONCAT_INNER(x, y)
#define _CG_NAMESPACE _CG_CONCAT_OUTER(__v, _CG_ABI_VERSION)

#define _CG_BEGIN_NAMESPACE      \
    namespace cooperative_groups \
    {                            \
        namespace _CG_NAMESPACE  \
        {
#define _CG_END_NAMESPACE          \
    }                              \
    ;                              \
    using namespace _CG_NAMESPACE; \
    }                              \
    ;

#if (defined(__cplusplus) && (__cplusplus >= 201103L)) || (defined(_MSC_VER) && (_MSC_VER >= 1900))
#define _CG_CPP11_FEATURES
#endif

#if !defined(_CG_QUALIFIER)
#define _CG_QUALIFIER __forceinline__ __device__
#endif
#if !defined(_CG_STATIC_QUALIFIER)
#define _CG_STATIC_QUALIFIER static __forceinline__ __device__
#endif
#if !defined(_CG_CONSTEXPR_QUALIFIER)
#if defined(_CG_CPP11_FEATURES)
#define _CG_CONSTEXPR_QUALIFIER constexpr __forceinline__ __device__
#else
#define _CG_CONSTEXPR_QUALIFIER _CG_QUALIFIER
#endif
#endif
#if !defined(_CG_STATIC_CONSTEXPR_QUALIFIER)
#if defined(_CG_CPP11_FEATURES)
#define _CG_STATIC_CONSTEXPR_QUALIFIER static constexpr __forceinline__ __device__
#else
#define _CG_STATIC_CONSTEXPR_QUALIFIER _CG_STATIC_QUALIFIER
#endif
#endif

#if defined(_MSC_VER)
#define _CG_DEPRECATED __declspec(deprecated)
#else
#define _CG_DEPRECATED __attribute__((deprecated))
#endif

#if defined(__MACA_ARCH__)
#define _CG_HAS_GRID_GROUP
#endif
#if defined(__MACA_ARCH__)
#define _CG_HAS_MULTI_GRID_GROUP
#endif
#if defined(__MACA_ARCH__)
#define _CG_HAS_MATCH_COLLECTIVE
#endif

// Has __half and __half2
// Only usable if you include the cuda_fp16.h extension, and
// _before_ including cooperative_groups.h
#ifdef __MACA_FP16_TYPES_EXIST__
#define _CG_HAS_FP16_COLLECTIVE
#endif

#if defined(__MACA_ARCH__)
#define _CG_HAS_OP_REDUX
#endif

#if defined(CG_USE_MACA_STL)
#define _CG_USE_MACA_STL
#else
#define _CG_USE_OWN_TRAITS
#endif

#ifdef _CG_CPP11_FEATURES
// Use cuda::std:: for type_traits
#if defined(_CG_USE_MACA_STL)
#define _CG_STL_NAMESPACE cuda::std
#include <cuda/std/type_traits>
// Use CG's implementation of type traits
#else
#define _CG_STL_NAMESPACE cooperative_groups::details::templates
#endif
#endif

#ifdef _CG_CPP11_FEATURES
#define _CG_STATIC_CONST_DECL static constexpr
#define _CG_CONST_DECL constexpr
#else
#define _CG_STATIC_CONST_DECL static const
#define _CG_CONST_DECL const
#endif

#if (defined(_MSC_VER) && !defined(_WIN64)) || defined(__arm__)
#define _CG_ASM_PTR_CONSTRAINT "r"
#else
#define _CG_ASM_PTR_CONSTRAINT "l"
#endif

/*
** Define: CG_DEBUG
** What: Enables various runtime safety checks
*/
#if defined(__MACACC_DEBUG__) && defined(CG_DEBUG) && !defined(NDEBUG)
#define _CG_DEBUG
#endif

#if defined(_CG_DEBUG)
#include <assert.h>
#define _CG_ASSERT(x) assert((x));
#define _CG_ABORT() assert(0);
#else
#define _CG_ASSERT(x)
#define _CG_ABORT() __trap();
#endif

#endif // _CG_INFO_H_